__attribute__((global, cooperative)) void reduce_sum_of_squares(float *inp, float *out, size_t elems_per_thread)
{
    // 定义共享内存
    __shared__ float sdata[12];
    // 全局索引
    auto thread_idx = threadIdx.z * (blockDim.x * blockDim.y) + (gridDim.x * threadIdx.y + blockIdx.x) * blockDim.x + threadIdx.x;
    // 定义DTE
    tops_dte_ctx_t dte_input[2];
    dte_input[0].init();
    dte_input[1].init();

    // 事件: 输入DMA完成
    tops::event event_input0;
    tops::event event_input1;

    static const size_t L1_BUFFER_LEN = 256;
    auto pos_start = thread_idx * elems_per_thread;
    auto pos_end = pos_start + elems_per_thread;

    __valigned__ float input_buffer[L1_BUFFER_LEN];

    tops::mdspan l1_input(tops::Private, &input_buffer, L1_BUFFER_LEN);
    tops::mdspan l1_out(tops::Private, &l1_buf_out, L1_BUFFER_LEN);

    for (size_t i = pos_start; i < pos_end; i += L1_BUFFER_LEN)
    {
        //! bind L3 Buffer to mdspan
        tops::mdspan l3_inp(tops::Global, inp + i, L1_BUFFER_LEN);
        //! memcpy from L3 to L1.
        tops::memcpy(dte_input[0], l1_input, l3_inp);
        for (size_t j = 0; j < L1_BUFFER_LEN; j += tops::vlength<vfloat>())
        {
            //! load data from L1 to vec register
            const auto &v_inp = tops::vload<vfloat>(input_buffer + j);
            //! do add
            const auto &v_res = tops::vmul<vfloat>(v_inp, v_inp);
            //! store data to L1
            tops::vstore(v_res, l1_buf_out + j);
        }
    }

    // 销毁DTE
    dte_input[0].destroy();
    dte_input[1].destroy();
}
void GCU_EUC_NORM(float *__restrict dev_inp, float *__restrict dev_out, const int nr_elems)
{

    static const size_t blocks = 1;
    static const size_t threads = 12;

    reduce_sum_of_squares<<<blocks, threads>>>(dev_inp, dev_out, nr_elems / (blocks * threads));
}